[CK_TILE] Allow UniversalGemmKernel::RunGemm to be called using tensor descriptors#3457
Draft
amd-meskelin wants to merge 4 commits intodevelopfrom
Draft
[CK_TILE] Allow UniversalGemmKernel::RunGemm to be called using tensor descriptors#3457amd-meskelin wants to merge 4 commits intodevelopfrom
UniversalGemmKernel::RunGemm to be called using tensor descriptors#3457amd-meskelin wants to merge 4 commits intodevelopfrom
Conversation
amd-meskelin
force-pushed
the
meskelin/add_desc_tensorviews_universal
branch
from
160bd91 to
ee9ba8c
Compare
7 tasks
|
Error importing due to merge conflicts – please reopen the PR on ROCm/rocm-libraries |
7 tasks
ThomasNing
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Feb 24, 2026
…ller to pass desciptors directly (#4295) ## Proposed changes Currently `UniversalGemmKernel::MakeA/B/C/DBlockViews` directly create tensor views from strides and sizes. This refactors the descriptor creation out and add overloaded definitions, allowing descriptors to be created separately by the caller instead of passing explicit strides, with no functional changes. This will enable further refactoring of `RunGemm` to do likewise, enabling derived kernels like BatchedContractionKernel to avoid creating separate versions (PR [#3457](ROCm/composable_kernel#3457)). ## Checklist Please put an `x` into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask. - [x] I have added tests relevant to the introduced functionality, and the unit tests are passing locally - [ ] I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, **IF** the test takes more than 30 seconds to run. - [ ] I have added inline documentation which enables the maintainers with understanding the motivation - [ ] I have removed the stale documentation which is no longer relevant after this pull request - [ ] (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request - [x] I have run `clang-format` on all changed files - [ ] Any dependent changes have been merged ## Discussion Since the logic within the MakeXBlockviews chains together operations on tuples, and thus the descriptors are also passed as such, adding a template parameter for the type of the input tuple was the simplest option to enable the overload without too much verbiage. However, for `MakeCBlockView` this adds a complications as the templated definitions are prone to overlap. This for now is avoided by just moving the arguments around for the descriptor version, which avoids the collision. It's not a great solution, so feel free to suggest a better one. --- 🔁 Imported from [ROCm/composable_kernel#3467](ROCm/composable_kernel#3467) 🧑💻 Originally authored by @amd-meskelin --------- Co-authored-by: Matti Eskelinen <matti.eskelinen@amd.com> Co-authored-by: systems-assistant[bot] <systems-assistant[bot]@users.noreply.github.com> Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com> Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
assistant-librarian bot
pushed a commit
that referenced
this pull request
Feb 24, 2026
[CK_TILE] Refactor `UniversalGemm::MakeA/B/C/DBlockViews` to allow caller to pass desciptors directly (#4295) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Proposed changes Currently `UniversalGemmKernel::MakeA/B/C/DBlockViews` directly create tensor views from strides and sizes. This refactors the descriptor creation out and add overloaded definitions, allowing descriptors to be created separately by the caller instead of passing explicit strides, with no functional changes. This will enable further refactoring of `RunGemm` to do likewise, enabling derived kernels like BatchedContractionKernel to avoid creating separate versions (PR [#3457](#3457)). ## Checklist Please put an `x` into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask. - [x] I have added tests relevant to the introduced functionality, and the unit tests are passing locally - [ ] I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, **IF** the test takes more than 30 seconds to run. - [ ] I have added inline documentation which enables the maintainers with understanding the motivation - [ ] I have removed the stale documentation which is no longer relevant after this pull request - [ ] (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request - [x] I have run `clang-format` on all changed files - [ ] Any dependent changes have been merged ## Discussion Since the logic within the MakeXBlockviews chains together operations on tuples, and thus the descriptors are also passed as such, adding a template parameter for the type of the input tuple was the simplest option to enable the overload without too much verbiage. However, for `MakeCBlockView` this adds a complications as the templated definitions are prone to overlap. This for now is avoided by just moving the arguments around for the descriptor version, which avoids the collision. It's not a great solution, so feel free to suggest a better one.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Proposed changes [WIP]
BatchedContractionKernelimplements a customRunGemmthat can be called with tensor descriptors, unlikeUniversalGemmKernel::RunGemmwhich computes tensor views directly from the given stride parameters.This adds
UniversalGemmKernel::RunGemmDescwhich takes in descriptors and thus replaces the custom implementation inBatchedContractionKernel. It also has a (partial) POC to implement the originalUniversalGemmKernel::RunGemmusing the descriptor version, see discussion below.Checklist
Please put an
xinto the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.clang-formaton all changed filesDiscussion
UniversalGemmKernel::RunGemmgenerates tensor views directly from the given parameters throughMakeGemmTensorViews, rather then sepaately building the descriptors and then generating views. In order to refactor the rest of the code path toRunGemmDesc(which would make it simple to have both entry points available),MakeGemmTensorViewwould need to be refactored to separately build the descriptors and views.OTOH,
GroupedConvolutionKernelalready has a similarMakeGemmTensorViewwhich takes in descriptors as doesRunGemm, with no version that expects to construct these internally. ShouldUniversalGemmKernelfollow suite? Alternately, could we expect the caller to construct & pass in the tensor views directly, instead of just the descriptors?